Doing a

        idx = detail::compress_a_impl<ScalarT, CompressedSize>(v);
        return *static_cast<const VecCompressed*>(&v);

here gives us an error:

sparse_transforms.hpp:79:17: error: static_cast from '_Float16 * __attribute__((ext_vector_type(16)))' to 'const VecCompressed *' (aka 'const ck_tile::impl::ext_vector<_Float16, 8>::type *') is not allowed
   79 |         return *static_cast<const VecCompressed*>(&v);

Can we avoid copying ? Ideas?

the buffer earlier used reinterpret_cast, maybe it works here?

Another option is to use what the original code did:

const AVecCompressed a_vec_pruned = {a_vec[0], a_vec[1], a_vec[2], a_vec[3]};

With some if constexpr or some sort of template for different sizes. There may be some perf reason why it was implemented like this.

reinterpret_cast gives a similar error, unfortunately.
What the "original" code did is still creating a new vector and copying over elements to it. It is also not universal since the compressed A vector doesn't have the same size for GFX9 and GFX12...
I will investigate a way to change the size of the ext_vector without unnecessary copies.

I see... it looks like it's gonna be hard to deal with with w/o creating a copy. If there is nothing we can do, the original code will be the safest option to go with

See my comment here: #5508 (comment)

This change is not good. It was done because DEVICE and HOST code are weirdly intertwined and I'll think of a way to revert this.

I'm not sure if all these operations belong here and whether we should treat them the same as compress_a. Did we discuss this already? I thought *ATTR_NUM_ACCESS can always be computed deterministically, i.e. not a pipeline option

Yeah I think Swizzle and AttrNumAccess are just modifiers on the tile distribution encoding and do not affect the vector fragments. Same for CTranspose except the vector fragments get swapped. I don't know if there is a need for explicit transposition anywhere.

Well the CTranspose can be kept as a pipeline implementation in the base class. Thoughts?

the buffer earlier used reinterpret_cast, maybe it works here?

Do we want this to exist separately from the existing compress_a function?

Well the idea is that the compress_a function will be removed when we integrate this to CK Tile. I also reverted the changes done to that function to minimize untested regression.

CK library generally uses remove_cvref_t<> instead.

I didn't use that as it is C++20, didn't know ck has its own implementation. I'll use that, thanks!

Very verbose but I guess necessary if we really don't want to allow raw enums...

Aha we are making a second-order wrapper for the intrinsic just like in CK Tile, making more sense to me now.

See also this thread of comments: #5508 (comment)

I used reinterpret_cast with std::launder. TBF this might have implications that I (and we) cannot foresee. You can't simply remove UB with an std::launder. Or maybe I'm wrong and that's the correct way to go for C++17. I'll research the internals of std::bit_cast to make an educated guess.
This code should definitely be removed when we stop caring about C++17.

For the time being this segment is tested and works as expected.

hmm, isn't this always true?

this works only if A and B have the same type and size

Yeah, CTranspose will not be available for all intrinsics. Also I don't think CTranspose is possible for sparse intrinsics.

I'll disable them for sparse then!

I think we should be very clear about the distinction between MmaOp flags and Pipeline flags. Also maybe we can come up with a better name than "pipeline", because pipeline already exists as a term for the wavegroup-sized full pipeline 1-2 levels higher.

Without giving much thought on it I am thinking MmaComposition but composition makes me think of using different mma structs. So it is not really a composition.
If you guys have any ideas please reach out!